Drop cub::ThreadReduce from the device sort (fixes 0.0.9 render crash) - #65
Merged
Merged
Conversation
duburcqa
force-pushed
the
fix-cccl-bundling
branch
2 times, most recently
from
July 23, 2026 06:35
0d71006 to
a5a3681
Compare
duburcqa
changed the title
The batch renderer JIT-compiles its device code with nvrtc at first render. sort_archetype.cpp summed a per-part histogram with cub::ThreadReduce, but that symbol moved namespaces across CUB versions (cub::internal:: in <=2.7, cub:: in 2.8+), and Genesis-Embodied-AI#61 rewrote the call to the unqualified 2.8 form. The release wheel bundles CUDA 12.8's CUB (2.7.0), so nvrtc failed with "identifier ThreadReduce is undefined" (NVRTC_ERROR_COMPILATION) -> SIGABRT, breaking every batch-render test in 0.0.9. The reduction just sums the NUM_PARTS partial histograms (NUM_PARTS == 1 in practice, so the branch is never taken at runtime). Replace ThreadReduce with a plain loop, removing the CUB-version dependency entirely. Also pin the wheel build to the 12.8 toolkit (/usr/local/cuda symlink + explicit cuda-cccl-12-8) for a consistent CUDA version, and print the bundled CUB version from CMake to surface future header skew.
duburcqa
force-pushed
the
fix-cccl-bundling
branch
from
July 23, 2026 06:53
a5a3681 to
57f2ddb
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Problem
0.0.9's batch renderer aborts at first render: nvrtc fails to compile the bundled device code with
identifier "ThreadReduce" is undefined(NVRTC_ERROR_COMPILATION→ SIGABRT), crashing all 37 batch-render tests.Cause:
sort_archetype.cppreduces a per-part histogram withcub::ThreadReduce. That symbol lives atcub::internal::in CUB ≤ 2.7 and atcub::in 2.8+. #61 rewrote the call to the unqualified 2.8 form (ThreadReduce(...)viausing namespace cub;), but the release wheel bundles CUDA 12.8's CUB (2.7.0), where the unqualified name doesn't resolve. (My #61 validation only used editable cluster builds against CUDA 12.9 → CUB 2.8.2, so the broken wheel went unnoticed until 0.0.9 was render-tested.)Fix
The reduction just sums the
NUM_PARTSpartial histograms, andNUM_PARTS == 1in practice (so the block is never even taken at runtime). ReplaceThreadReducewith a plain unrolled loop — removing the CUB-version dependency entirely (nointernal::guessing, no CUDA-version bump; stays on 12.8).Also, for build hygiene:
/usr/local/cuda→ the 12.8 toolkit inbefore-all(+ explicitcuda-cccl-12-8) so the build resolves one consistent CUDA version instead of a pre-existing symlink.Validation
The device code is JIT-compiled at runtime, so the wheel-build CI only confirms the wheel compiles — it can't render. Validation is via a dev pre-release render-tested through genesis (Production job on the cluster) before cutting a real patch.